#include "kernel_operator.h"
#include "kernel_select_fp32.h"
#include "kernel_select_fp16.h"
#include "kernel_select_int32.h"
#include "kernel_select_int8.h"
#include "kernel_select_broadcast.h"
extern "C" __global__ __aicore__ void select_v2(GM_ADDR condition, GM_ADDR x1, GM_ADDR x2, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)
{
    TPipe pipe;
    GET_TILING_DATA(tiling_data, tiling);
    if (TILING_KEY_IS(5)){
        KernelBroadCast<uint8_t, float> op;
        op.Init(condition, x1, x2, y,
                tiling_data.y_dimensional,
                tiling_data.y_ndarray, tiling_data.x0_ndarray, tiling_data.x1_ndarray, tiling_data.x2_ndarray,
                tiling_data.y_sumndarray, tiling_data.x0_sumndarray, tiling_data.x1_sumndarray, tiling_data.x2_sumndarray);
        op.Process();
    }else if (TILING_KEY_IS(1)){
        KernelFp32<uint8_t, float> op;
        op.Init(&pipe, condition, x1, x2, y,
                tiling_data.coreDataNum, tiling_data.finalTileNum, tiling_data.tileDataNum, tiling_data.tailDataNum);
        op.Process();
    }else if (TILING_KEY_IS(2)){
        KernelFp16<uint8_t, half> op;
        op.Init(&pipe, condition, x1, x2, y,
                tiling_data.coreDataNum, tiling_data.finalTileNum, tiling_data.tileDataNum, tiling_data.tailDataNum);
        op.Process();
    }else if (TILING_KEY_IS(3)){
        KernelInt32<uint8_t, int32_t> op;
        op.Init(&pipe, condition, x1, x2, y,
                tiling_data.coreDataNum, tiling_data.finalTileNum, tiling_data.tileDataNum, tiling_data.tailDataNum);
        op.Process();
    }else if (TILING_KEY_IS(4)){
        KernelInt8<uint8_t, int8_t> op;
        op.Init(&pipe, condition, x1, x2, y,
                tiling_data.coreDataNum, tiling_data.finalTileNum, tiling_data.tileDataNum, tiling_data.tailDataNum);
        op.Process();
    }
}
